//----------------------------------------------------------------------------
// M A R I T I M E  R E S E A R C H  I N S T I T U T E  N E T H E R L A N D S
//----------------------------------------------------------------------------
// Copyright (C) 2012 - MARIN - All rights reserved - http://www.marin.nl
//----------------------------------------------------------------------------
// Program    : mermaid
// Module     : vbm
// File       : dotp.h
// Author     : M.A. de Jong
//----------------------------------------------------------------------------

#ifndef DOTP_H
#define	DOTP_H

#include "kernels/dotp_kernels.h"

namespace rrbsolver
{

template <class T>
inline void dotp_phase1(T *hrem, T *drem, const T *dY, const T *dX, const int size,
    const Grid grid)
{
    dimGrid.x = grid.cx / DIM_COMPUTE_BLOCK;
    dimGrid.y = grid.cy / DIM_COMPUTE_BLOCK;
    dimBlock.x = DIM_COMPUTE_BLOCK;
    dimBlock.y = DIM_COMPUTE_BLOCK / DOTP_BF;

    kernel_dotp1<T> <<< dimGrid, dimBlock >>> (drem, dY, dX, grid);
    checkCudaError(cudaGetLastError(), "rrbsolver::kernel_dotp1");

    checkCudaError(cudaMemcpy(hrem, drem, size * sizeof(T), 
        cudaMemcpyDeviceToHost), "rrbsolver::memcpy remainder");
}

} // namespace rrbsolver

#endif // DOTP_H

